…L declaration

Co-authored-by: ssheorey <41028320+ssheorey@users.noreply.github.com>

…l-kernels-for-cuda

Some optimizations. (e.g. restrict)
lock free hashmap insert: write buffer then CAS design is correct and fast, but leaves holes in the data buffer.
Center data before Knn, if using expanded L2 distance formula (p^2+q^2-2pq) to prevent cancellation
bug in indexer: TensorIterator::GetPtr() incorrect for non-contiguous.
Added Knn search benchmark

…faster!

- SYCLContext: process-wide static singleton (was thread_local) to avoid
  per-thread SYCL contexts/USM mismatches; cache all device properties in
  one place (SYCLContext::Impl) and expose via GetDeviceProperties().
- Add SYCL launch helpers (SYCLPreferredWorkGroupSize, SYCLNdRange1D,
  group-reduction helpers) and use nd_range<1> + sycl::reduce_over_group
  across ParallelFor, elementwise, and reduction kernels instead of flat
  parallel_for(n) / per-output kernel launches.
- ReductionSYCL: on-device GetInputPtrDevice() enables one kernel with one
  work-group per output for multi-output reductions (incl. arg-reductions).
- Registration/RGBD odometry/SLAC kernels: accumulate AtA/Atb/residual in
  SLM per work-group (restrict-qualified pointers) instead of global atomics.
- Rename SYCLBlockCopyDispatch.h -> BlockCopyDispatch.h, generalize the
  vectorized object-copy dispatch (up to 64-byte blocks) shared by the hash
  map and tensor copy paths; align CUDA hashmap Dispatch.h divisors to match.
- Build: set -fsycl-max-parallel-link-jobs, prefer lld linker when available.
- Update ParallelFor/Reduction benchmarks and Linalg/Tensor tests for the
  refactored SYCLContext API.

Add core::sy::IsCPUDevice() to detect the SYCL CPU fallback device and
use it to throw clear errors from LeastSquaresSYCL (gels_batch), the
SYCL hash map, and FixedRadiusSearch/HybridSearch, which are broken on
SYCL CPU. Skip the corresponding C++ tests (HashMap, NNS, VoxelBlockGrid,
Registration, Feature, PointCloud) and opt affected Python tests out of
the SYCL CPU fallback via list_devices(also_sycl_cpu=False).

- SYCL CPU fallback: Restored IsCPUDevice guards for FixedRadiusSearch
  and HybridSearch (and their dependent modules: NearestNeighborSearch,
  PointCloud, Registration, Feature) due to large-scale segmentation
  faults on SYCL CPU. HashMap and VoxelBlockGrid are now fully supported
  on SYCL CPU.
- MSVC preprocessor: Fixed compilation error in PointCloudImpl.h by
  replacing nested #if/#elif inside lambda with OPEN3D_ATOMIC_ADD.
- CUDA compilation: Resolved BlockCopy64 template argument error in
  SlabHashBackend/StdGPUHashBackend by moving struct to namespace scope,
  and brought global isnan into trianglemesh namespace.
- Benchmarks: Explicitly linked Open3D::3rdparty_sycl to benchmarks
  target to resolve DSO link errors.
- Tests: Updated Python and C++ test suites to run cleanly with
  ONEAPI_DEVICE_SELECTOR=opencl:cpu and when SYCL is disabled.

Co-authored-by: Cursor <cursoragent@cursor.com>

… make CMake link options private

This prevents runtime crashes on CPU devices, improves performance by caching contiguity, and avoids propagating build-time options to downstream consumers.

Co-authored-by: Cursor <cursoragent@cursor.com>

This cleans up the PR's diff and ensures compliance with the CI's style-checking rules.

Co-authored-by: Cursor <cursoragent@cursor.com>

…brid NNS

Replace the O(N*M) tiled-AddMM brute force in FixedRadiusSearchSYCL and
HybridSearchSYCL with a shared spatial-hash CSR grid (cell size 2*radius),
mirroring CUDA's FixedRadiusSearchImpl.cuh: count -> oneDPL inclusive_scan
-> scatter to build the grid, then a two-pass count/gather (radius) or
single-pass top-max_knn (hybrid) over the 8 corner-adjacent hash bins per
query. sort=true uses an on-device oneDPL sort_by_key (packed uint64 radix
key for float, struct key + comparator for double).

This also removes the SYCL-CPU restriction for FixedRadius/Hybrid search,
since the grid algorithm (unlike the old AddMM path) has no GPU-specific
sizing dependency; verified by unskipping and passing the previously
CPU-guarded C++/Python parity tests, plus EstimateNormals/
RemoveRadiusOutliers end-to-end on SYCL:0 CPU fallback.

Co-authored-by: Cursor <cursoragent@cursor.com>

…ign note

Fixes the style-check CI failure (line-wrapping only, no logic changes).
Also adds docs/cuda_vs_sycl_nns_hashmap.md summarizing how the SYCL and
CUDA core::nns/core::HashMap implementations compare after the
fixed-radius/hybrid grid port.

Co-authored-by: Cursor <cursoragent@cursor.com>

Pre-reserve one heap slot per insert thread on the host (CUDA SlabHash-style)
and DeviceFree slots that lose the CAS or hit a duplicate; avoids in-kernel Allocate/Free races.

RaycastingScene: dedupe ray-triangle hits by (geomID, tfar) instead of
(geomID, primID, tfar), since Embree's SYCL/GPU filter callback does not
reliably refresh primID across successive candidate hits for one ray,
causing genuinely distinct intersections to be dropped.

ParallelFor: guard the non-vectorized overload with device.IsSYCL() before
dispatching to ParallelForSYCL_, matching the vectorized overload, so CPU
devices don't crash when called from a SYCL-compiled translation unit.

Also includes prior fixes from this session: SYCL reduction_prod tree
reduction, KnnSearchSYCL large-k distance clamping, a TensorFunction
tail-call stack-unwinding crash, and enabling SYCL devices by default in
Python test parametrization.

Co-authored-by: Cursor <cursoragent@cursor.com>

…l-kernels-for-cuda

…ug fix.

Consolidate shared registration/transform/feature logic, SYCLUtils helpers,
and hashmap buf_indices handling; expand in-source NNS/hash documentation and
refresh C++/Python SYCL tests (drop redundant voxel block grid Python test).

Co-authored-by: Cursor <cursoragent@cursor.com>

Guard ComputeRtPointToPointTensor when neither CUDA nor SYCL is built; apply
style; skip SYCL-only C4 tie-break on CUDA; defer Windows SYCL VoxelBlockGrid
golden tests; add CHANGELOG entry for SYCL tensor backends.

Co-authored-by: Cursor <cursoragent@cursor.com>

Co-authored-by: Cursor <cursoragent@cursor.com>